#include <stdio.h>
#include <stdlib.h>

// includes, project
#include "cutil_inline.h"
//#include <cuda.h>
//#include <shrQATest.h>
typedef float2 Complexgf;

__global__ void imageCopyKernal(unsigned char* src,unsigned char* srcOneBefore,unsigned char* srcTwoBefore, unsigned char* dest,int width,int height,int noOfImages) 
{
    // Each thread computes one element of C 
	// by accumulating results into Cvalue 

	float Cvalue = 0; 
	const float lpf[2]={-1, 1};
 
	int row = blockIdx.y * blockDim.y + threadIdx.y; 
	int col = blockIdx.x * blockDim.x + threadIdx.x; 
	
	//int row = (blockIdx.z-1)*blockDim.x*blockDim.y  + blockIdx.y * blockDim.y + threadIdx.y;  // check the -1
	//int col =(blockIdx.z-1)*blockDim.x*blockDim.y+ blockIdx.x * blockDim.x + threadIdx.x; 

	//int col = threadIdx.x; 
	//int row = blockIdx.x;

   /* for (int e = 0; e < 2; ++e) 
		Cvalue += src[blockIdx.z*width*height +row * width + col+ e]*lpf[e]; 
	
	dest[blockIdx.z*width*height+row * width + col+1] = Cvalue;*/
	int diff = abs(src[blockIdx.z*width*height +row * width + col]- srcOneBefore[blockIdx.z*width*height +row * width + col]);
	//dest[blockIdx.z*width*height+row * width + col]=srcTwoBefore[blockIdx.z*width*height +row * width + col];
   // dest[blockIdx.z*width*height+row * width + col] = 0; 
    if(diff>=3)
    {
       // dest[blockIdx.z*width*height+row * width + col]=abs(src[blockIdx.z*width*height +row * width + col]- srcOneBefore[blockIdx.z*width*height +row * width + col])+abs(srcOneBefore[blockIdx.z*width*height +row * width + col]-srcTwoBefore[blockIdx.z*width*height +row * width + col]);
       // dest[blockIdx.z*width*height+row * width + col]=abs(src[blockIdx.z*width*height +row * width + col]-srcOneBefore[blockIdx.z*width*height +row * width + col]);
        dest[blockIdx.z*width*height+row * width + col] = diff;
        //dest[blockIdx.z*width*height+ col]= dest[blockIdx.z*width*height+ col]+diff;
        //dest[blockIdx.z*width*height+row * width] = dest[blockIdx.z*width*height+row * width] + diff;
        
    }
    
    //src[blockIdx.z*width*height +row * width + col]-
  // use different diamentions for the different cameras
}

__global__ void maskKernal(unsigned char* src,unsigned char* dest,int width,int height,int noOfImages) 
{
    float Cvalue = 0; 
	int row = blockIdx.y * blockDim.y + threadIdx.y; 
	int col = blockIdx.x * blockDim.x + threadIdx.x; 
	for (int e = 0; e < width; ++e)
	{
	    Cvalue+= src[blockIdx.z*width*height+row * width + e];
	}
	
	dest[blockIdx.z*width*height+row * width] = Cvalue;
	

}


extern "C" bool imageCopy(unsigned char* src,unsigned char* srcOneBefore,unsigned char* srcTwoBefore,unsigned char* dest,int width,int height,size_t size,int noOfImages)
{
      //  int i=16;
        int blockSize=32;
        
        int deviceCount = 0;

        cudaGetDeviceCount(&deviceCount);
        if (deviceCount == 0)
        {
                printf("There are no available device(s) that support CUDA\n");
        }
        unsigned char* d_data;
        unsigned char* d_dataOneBefore;
        unsigned char* d_dataTwoBefore;
        
        unsigned char* d_data_temp;
         unsigned char* d_data_out;
         
        cudaMalloc(&d_data, size);
        cudaMalloc(&d_dataOneBefore, size);
        cudaMalloc(&d_dataTwoBefore, size);
        
                       
        cudaMemcpy(d_data, src, size, cudaMemcpyHostToDevice);
        cudaMemcpy(d_dataOneBefore, srcOneBefore, size, cudaMemcpyHostToDevice);
        cudaMemcpy(d_dataTwoBefore, srcTwoBefore, size, cudaMemcpyHostToDevice);
        
       /* cudaFree(d_dataTwoBefore);
        cudaMalloc(&d_dataTwoBefore, size);
        cudaMemcpy(d_dataTwoBefore, d_dataOneBefore, size, cudaMemcpyDeviceToDevice);
        cudaFree(d_dataOneBefore);        
        cudaMalloc(&d_dataOneBefore, size);
        cudaMemcpy(d_dataOneBefore, d_data, size, cudaMemcpyDeviceToDevice);*/
        
        //temporary image data
	    cudaMalloc(&d_data_temp, size);
	    cudaMalloc(&d_data_out, size);
        
        //// Invoke kernel X 
	    dim3 dimBlock(blockSize,blockSize); 
	    dim3 dimGrid(ceil(width/(float)blockSize),ceil(height/(float)blockSize),noOfImages);
        imageCopyKernal<<<dimGrid, dimBlock>>>(d_data,d_dataOneBefore,d_dataTwoBefore, d_data_temp,width,height,noOfImages);
        cudaThreadSynchronize();
        
        //maskKernal<<<dimGrid, dimBlock>>>(d_data_temp,d_data_out,width,height,noOfImages);
    //   __global__ void d_recursiveGaussianY(int *d_src, int *d_dest, int depth, int height, int width, float b0, float b1, float b2, float b3, float B, int order, float M11, float M12, float M13, float M21, float M22, float M23, float M31, float M32, float M33)
        
        /*cudaFree(d_dataTwoBefore);
        cudaMalloc(&d_dataTwoBefore, size);
        cudaMemcpy(d_dataTwoBefore, d_dataOneBefore, size, cudaMemcpyDeviceToDevice);
        cudaFree(d_dataOneBefore);        
        cudaMalloc(&d_dataOneBefore, size);
        cudaMemcpy(d_dataOneBefore, d_data, size, cudaMemcpyDeviceToDevice);*/
         
        
        // Read the output from the device memory 
	    //cudaMemcpy(dest, d_data_temp, size, cudaMemcpyDeviceToHost); 
	    cudaMemcpy(dest, d_data_out, size, cudaMemcpyDeviceToHost); 
        
        
        cudaFree(d_data);
        cudaFree(d_dataOneBefore);
        cudaFree(d_dataTwoBefore);
        cudaFree(d_data_temp);

        return true;

}
 __global__ void gaborDistanceKernal(float* d_data_energy,Complexgf* d_resultPrevious,Complexgf* d_resultNow, int width,int height) 
{    

    /*unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;     
    unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int Index = yIndex * width + xIndex;*/
    
   /* unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;*/
    
    /*int row = blockIdx.x * blockDim.x + threadIdx.x; 
	int col = blockIdx.y * blockDim.y + threadIdx.y;*/
    
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
	int col = blockIdx.x * blockDim.x + threadIdx.x;
	
	/*if(row==0&&col==0)
	{
	
	    printf("\nDistance :%d,%d\n",row,height);
	
	} */ 
	
	if (row < height && col < width) 
	{  	
	     // if(
	
	    float diffX = d_resultPrevious[row * width + col].x - d_resultNow[row * width + col].x;
	   // float diffX = d_resultNow[row * width + col].x - d_resultPrevious[row * width + col].x;
	    //float diffX = 10 - 10;
	    float diffY = d_resultPrevious[row * width + col].y - d_resultNow[row * width + col].y;      //row * width + col
	   // float diffY = d_resultNow[row * width + col].y - d_resultPrevious[row * width + col].y;      
	    //float diffY = 20-20;  
	  //   printf("Previous :%.2f\n",d_resultPrevious[row * width + col].y);
	   // if(row==10&&col==10)
	   //float ePrevious =  d_resultPrevious[row * width + col].x*d_resultPrevious[row * width + col].x+d_resultPrevious[row * width + col].y*d_resultPrevious[row * width + col].y;
	  // float eNow =  d_resultNow[row * width + col].x*d_resultNow[row * width + col].x+d_resultNow[row * width + col].y*d_resultNow[row * width + col].y;
	   float tanPrevious = d_resultPrevious[row * width + col].y/d_resultPrevious[row * width + col].x;
	   float tanNow = d_resultNow[row * width + col].y/d_resultNow[row * width + col].x;
	   
	  // float phaseShift=atan(d_resultPrevious[row * width + col].y/d_resultPrevious[row * width + col].x)- atan(d_resultNow[row * width + col].y/d_resultNow[row * width + col].x);
	   //float disparity=atan(((d_resultPrevious[row * width + col].y*d_resultNow[row * width + col].x)-(d_resultNow[row * width + col].y*d_resultPrevious[row * width + col].x))/(d_resultNow[row * width + col].x*d_resultPrevious[row * width + col].x+d_resultNow[row * width + col].y*d_resultPrevious[row * width + col].y));//- atan(d_resultNow[row * width + col].y/d_resultNow[row * width + col].x);
	   float disparity=atan((tanPrevious-tanNow)/(1+tanPrevious*tanNow));
	   
      /* if(diffY<0)
       {
         // diffY=diffY*-1;
       }
	    {
	      //printf("\nXdiff :%.2f, Ydiff :%.2f, Distance :%.2f\n",diffX,diffY,sqrtf(diffX*diffX+diffY*diffY));
	    }   */
	    //d_data_energy[row * width + col] = -1*(atan(d_resultPrevious[row * width + col].y/d_resultPrevious[row * width + col].x)- atan(d_resultNow[row * width + col].y/d_resultNow[row * width + col].x));
	  //  d_data_energy[row * width + col] =d_resultPrevious[row * width + col].x - d_resultNow[row * width + col].x;
	    // d_data_energy[row * width + col] =d_resultNow[row * width + col].y - d_resultPrevious[row * width + col].y;
	  //  printf("%\nDistance :%.2f\n",d_data_energy[row * width + col] );
	   d_data_energy[row * width + col]=0;
    //  if(  phaseShift<0)    // diffX<0&& diffY<0
	  if(  disparity>0)    // diffX<0&& diffY<0
	  {
	       // d_data_energy[row * width + col] = -1*sqrtf(d_resultNow[row * width + col].y*d_resultNow[row * width + col].y+d_resultNow[row * width + col].x*d_resultNow[row * width + col].x)+sqrtf(d_resultPrevious[row * width + col].y*d_resultPrevious[row * width + col].y+d_resultPrevious[row * width + col].x*d_resultPrevious[row * width + col].x);
	       //d_data_energy[row * width + col]=-1*d_resultNow[row * width + col].y;
	      d_data_energy[row * width + col] = diffX*diffX+diffY*diffY;    //sqrtf(    diffX*diffX+diffY*diffY
	      //d_data_energy[row * width + col] = ePrevious-eNow;  
	    //d_data_energy[row * width + col] =  phaseShift;
	  }
	   // d_data_energy[row * width + col] = diffX+diffY;    //sqrtf(    diffX*diffX+diffY*diffY
	   // d_data_energy[row * width + col] = diffY*diffX;    //sqrtf(diffY*diffY) sqrtf(    diffX*diffX+diffY*diffY    48 14   39 5
	 /* if(( diffX*diffX+diffY*diffY)>0)
	  {
	  printf("%\nDistance :%.2f\n",(diffX*diffX+diffY*diffY));
	  }*/
	}
	
	
	//d_data_energy[row * width + col] =  diffX*diffX+diffY*diffY;// sqrtf(diffX*diffX+diffY*diffY);
	
	
}
extern "C" bool gaborDistance(float *d_data_energy,Complexgf *d_resultPrevious,Complexgf *d_resultNow,int width,int height)
{
        int blockSize=32;
        
        int deviceCount = 0;

        cudaGetDeviceCount(&deviceCount);
        if (deviceCount == 0)
        {
                printf("There are no available device(s) that support CUDA\n");
        }   

        dim3 dimBlock(blockSize,blockSize); 
	    dim3 dimGrid(ceil(width/(float)blockSize),ceil(height/(float)blockSize));
	   // dim3 dimGrid(ceil(height/(float)blockSize),ceil(width/(float)blockSize));
        gaborDistanceKernal<<<dimGrid, dimBlock>>>(d_data_energy,d_resultPrevious,d_resultNow,width,height);
        cudaThreadSynchronize();
        
        return true;

}

__global__ void frontMovingSegmentKernal(float* d_data_energy,int* d_minColLocation,int* d_maxColLocation,int width,int height) 
{
    int row = blockIdx.x * blockDim.x + threadIdx.x; 
	//int col = blockIdx.x * blockDim.x + threadIdx.x;
	
	//if(row==500)
	{
	
	   // printf("\Row :%d, Col :\n",row);
	
	}	
	
	if(row < height)// && col < width) 
	{
	     d_minColLocation[row]=-1;
	     d_maxColLocation[row]=-1;
	     for(int j=0;j<width;j++)
	     {
	        if(d_data_energy[row*width+j]>1)
	        {
	            if(d_minColLocation[row]==-1)
	            {
	                d_minColLocation[row]=j;
	               // printf("from GPU :%d",row);
	            }
	            
	            d_maxColLocation[row]=j;
	           // break;
	        }  	     
	     }
	}
}

extern "C" bool frontMovingSegment(float* d_data_energy,int* d_minColLocation,int* d_maxColLocation,int width,int height) 
{
        int blockSize=32;
        
        int deviceCount = 0;

        cudaGetDeviceCount(&deviceCount);
        if (deviceCount == 0)
        {
                printf("There are no available device(s) that support CUDA\n");
        }   

       // dim3 dimBlock(blockSize,blockSize); 
        dim3 dimBlock(blockSize); 
	    dim3 dimGrid(ceil(height/(float)blockSize));
	   // dim3 dimGrid(ceil(height/(float)blockSize),ceil(width/(float)blockSize));
        frontMovingSegmentKernal<<<dimGrid, dimBlock>>>(d_data_energy,d_minColLocation,d_maxColLocation,width,height);
        cudaThreadSynchronize();          
        return true; 
}
 __global__ void imageCharToFloatKernal(float *d_imageDataFloat,char *d_imageDataChar,int widthStep,int width,int height)
{
     int row = blockIdx.y * blockDim.y + threadIdx.y; 
	 int col = blockIdx.x * blockDim.x + threadIdx.x;
	 
	if (row < height && col < width) 
	{
	    d_imageDataFloat[row * width + col]=(float)(unsigned char)d_imageDataChar[row * widthStep + col];
	}
  
  
}
extern "C" bool imageCharToFloat(float *d_imageDataFloat,char *d_imageDataChar,int widthStep,int width,int height)
{
        int blockSize=32;
        
        int deviceCount = 0;  
        cudaGetDeviceCount(&deviceCount);
        if (deviceCount == 0)
        {
                printf("There are no available device(s) that support CUDA\n");
        }  
        dim3 dimBlock(blockSize,blockSize); 
	    dim3 dimGrid(ceil(width/(float)blockSize),ceil(height/(float)blockSize));
	    
	    imageCharToFloatKernal<<<dimGrid, dimBlock>>>(d_imageDataFloat,d_imageDataChar,widthStep,width,height);
        cudaThreadSynchronize(); 
	    return true; 

}

/*__global__ void d_recursiveGaussianY(int *d_src, int *d_dest, int depth, int height, int width, float b0, float b1, float b2, float b3, float B, int order, float M11, float M12, float M13, float M21, float M22, float M23, float M31, float M32, float M33)

{
float wP1 = 0.f, wP2 = 0.f, wP3 = 0.f; int y = 0; float outF1 = 0.f, outF2 = 0.f, outF3 = 0.f; unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; unsigned int yy = blockIdx.y*width*height; if(x > width) return;

d_src += x + yy;

d_dest += x + yy;

wP1 = (
float)*d_src/sqrt(B); wP2 = wP1; wP3 = wP1; switch (order)

{
case 0:

{
for(y=0;y<height;y++)

{
float xC = (float)*d_src; float wC = (float)(xC - b1*wP1 - b2*wP2 - b3*wP3)/b0;

*d_dest = (
int)wC;

d_src += width; d_dest += width;

wP3 = wP2; wP2 = wP1; wP1 = wC;

}

d_src -= width;

d_dest -= width;
float up = (float)*d_src/(1.0+b1+b2+b3); float vp = (float)up/(1.0+b1+b2+b3); float out = 0.f;

out = (
float)M11*(*d_dest-up) + M12*(*(d_dest-width)-up) + M13*(*(d_dest-2*width)-up)+vp;

outF1 = (
float)M21*(*d_dest-up) + M22*(*(d_dest-width)-up) + M23*(*(d_dest-2*width)-up)+vp;

outF2 = (
float)M31*(*d_dest-up) + M32*(*(d_dest-width)-up) + M33*(*(d_dest-2*width)-up)+vp;

out *= B; outF1 *= B; outF2 *= B;

outF3 = outF2; outF2 = outF1; outF1 = out;

*d_dest = (
int)out;

d_src -= width;

d_dest -= width;
for(y=height-1-1;y>=0;y--)

{
float wC = (float)*d_dest;

out = (
float)(B*wC - b1*outF1 - b2*outF2 - b3*outF3)/b0;

*d_dest = (
int)out;

d_src -= width; d_dest -= width;

outF3 = outF2; outF2 = outF1; outF1 = out;

}

}
break; case 1:

{
float xP1 = (float)*(d_src); float xF1 = (float)*(d_src + 1*width);

wP1 = (
float)(xF1 - xP1)/2.0*(1.0+b1+b2+b3); wP3 = wP2 = wP1; float wC = ((xF1- xP1)/2.0 - b1*wP1 - b2*wP1 - b3*wP1)/b0;

*d_dest = (
int)wC;

d_src += width; d_dest += width;

wP3 = wP2; wP2 = wP1; wP1 = wC;
for(y=0;y<height-1-1;y++)

{

xP1 = (
float)*(d_src - width);

xF1 = (
float)*(d_src + width);

wC = (
float)((xF1- xP1)/2.0 - b1*wP1 - b2*wP2 - b3*wP3)/b0;

*d_dest = (
int)wC;

d_src += width; d_dest += width;

wP3 = wP2; wP2 = wP1; wP1 = wC;

}

xP1 = (
float)*(d_src - width);

xF1 = (
float)*d_src;

wC = (
float)((xF1- xP1)/2.0 - b1*wP1 - b2*wP2 - b3*wP3)/b0;

*d_dest = (
int)wC;

wP3 = wP2; wP2 = wP1; wP1 = wC;
float up = (*d_src - *(d_src-width))/2.0*(1.0+b1+b2+b3); float vp = up/(1.0+b1+b2+b3); float out = 0.f;

out = (
float)M11*(*d_dest-up) + M12*(*(d_dest-width)-up) + M13*(*(d_dest-2*width)-up)+vp;

outF1 = (
float)M21*(*d_dest-up) + M22*(*(d_dest-width)-up) + M23*(*(d_dest-2*width)-up)+vp;

outF2 = (
float)M31*(*d_dest-up) + M32*(*(d_dest-width)-up) + M33*(*(d_dest-2*width)-up)+vp;

out *= B; outF1 *= B; outF2 *= B;

outF3 = outF2; outF2 = outF1; outF1 = out;

*d_dest = (
int)out;

d_src -= width;

d_dest -= width;
for(y=height-2;y>=0;y--)

{
float wC = (float)*d_dest;

out = (
float)(B*wC - b1*outF1 - b2*outF2 - b3*outF3)/b0;

*d_dest = (
int)out;

d_src -= width; d_dest -= width;

outF3 = outF2; outF2 = outF1; outF1 = out;

}

}
break; case 2:

{
float xP1 = (float)*d_src; float xC = (float)*d_src; float xF1 = (float)*(d_src+width);

wP1 = 0.0/(1.0+b1+b2+b3); wP3 = wP2 = wP1;
float wC = (float)((xF1 - 2*xC + xP1) - b1*wP1 - b2*wP2 - b3*wP3)/b0;

*d_dest = (
int)wC;

d_src += width; d_dest += width;

wP3 = wP2; wP2 = wP1; wP1 = wC;
for(y=0;y<height-2;y++)

{

xC = (
float)*d_src;

xP1 = (
float)*(d_src-width);

xF1 = (
float)*(d_src+width);

wC = (
float)(xF1 - 2*xC + xP1 - b1*wP1 - b2*wP2 - b3*wP3)/b0;

*d_dest = (
int)wC;

d_src += width; d_dest += width;

wP3 = wP2; wP2 = wP1; wP1 = wC;

}

xC = (
float)*d_src;

xP1 = (
float)*(d_src-width);

xF1 = (
float)*d_src;

wC = (
float)(xF1 - 2*xC + xP1 - b1*wP1 - b2*wP2 - b3*wP3)/b0;

*d_dest = (
int)wC;

wP3 = wP2; wP2 = wP1; wP1 = wC;
float up = 0; float vp = up/(1.0+b1+b2+b3); float out = 0.f;

out = (
float)M11*(*d_dest-up) + M12*(*(d_dest-width)-up) + M13*(*(d_dest-2*width)-up)+vp;

outF1 = (
float)M21*(*d_dest-up) + M22*(*(d_dest-width)-up) + M23*(*(d_dest-2*width)-up)+vp;

outF2 = (
float)M31*(*d_dest-up) + M32*(*(d_dest-width)-up) + M33*(*(d_dest-2*width)-up)+vp;

out *= B; outF1 *= B; outF2 *= B;

outF3 = outF2; outF2 = outF1; outF1 = out;

*d_dest = (
int)out;

d_src -= width;

d_dest -= width;
for(y=height-2;y>=0;y--)

{
float wC = (float)*d_dest;

out = (
float)(B*wC - b1*outF1 - b2*outF2 - b3*outF3)/b0;

*d_dest = (
int)out;

d_src -= width; d_dest -= width;

outF3 = outF2; outF2 = outF1; outF1 = out;

}

}
break;

}

}
*/
/*#include "stdafx.h"
#include "cuda.h"
#include <stdio.h>
#include <stdlib.h>
#include <vector>
#include <time.h>
#include <ctype.h>
#include <ctime>
#include <cstdio>
 
void initialize(int* elements,int width,int height);
void printElements(int* elements,int width,int height);


// Thread block size 
__global__ void GPU_filterKernal_X(int* src, int* dest,int width,int height);

using namespace std;



int GPU_filter(int* src, int* dest, size_t size,dim3 dimBlock,dim3 dimGrid,int width,int height)
{ 
	int *d_data;
	int *d_data_temp;
	

	
	cudaMalloc(&d_data, size);
	cudaMemcpy(d_data, src, size, cudaMemcpyHostToDevice);

	//temporary image data
	cudaMalloc(&d_data_temp, size);
	
	  
	//// Invoke kernel X 
	//dim3 dimBlock(31,1); 
	//dim3 dimGrid(32,1);

	GPU_filterKernal_X<<<dimGrid, dimBlock>>>(d_data, d_data_temp,width,height); 
	//__syncthreads();

	cudaThreadSynchronize();

		
	// Read the output from the device memory 
	cudaMemcpy(dest, d_data_temp, size, cudaMemcpyDeviceToHost); 

	cudaFree(d_data); 
	cudaFree(d_data_temp); 
	
	return 0;
}

__global__ void GPU_filterKernal_X(int* src, int* dest,int width,int height) 
{ 
	// Each thread computes one element of C 
	// by accumulating results into Cvalue 

	float Cvalue = 0; 
	const float lpf[2]={-1, 1};
 
	int row = blockIdx.y * blockDim.y + threadIdx.y; 
	int col = blockIdx.x * blockDim.x + threadIdx.x; 

	//int col = threadIdx.x; 
	//int row = blockIdx.x;

	for (int e = 0; e < 2; ++e) 
		Cvalue += src[row * width + col+ e]*lpf[e]; 
	
	dest[row * width + col+1] = Cvalue;
}

int main( void ) 
{

    printf("Object Tracking\n");

    
     clock_t start;
	//time_t start,end;
	double diff=0.0;

	int i=64;
	int blockSize=32;
	//for(int i=32;i<=4096;i++)
	{

	unsigned int width=i, height=i;
	size_t size = width * height * sizeof(int);

       int* h_data =(int*)malloc(size);  
	initialize(h_data,width,height);

	
	printElements(h_data,width,height);

	int deviceCount = 0;
 
	cudaGetDeviceCount(&deviceCount);
	if (deviceCount == 0) 
	{
        	printf("There are no available device(s) that support CUDA\n");
    	} 

	int* h_output = (int*) malloc( size);

	//// Invoke kernel X 
	dim3 dimBlock(blockSize,blockSize); 
	dim3 dimGrid(i/blockSize,i/blockSize);

	//dim3 dimBlock(256,256); 
	//dim3 dimGrid(4,4);


  	start = clock();	
	//printf("%f\n\n",(float)start);


	//time (&start);
	

	GPU_filter(h_data,h_output,size,dimBlock,dimGrid,width,height);
	//cudaThreadSynchronize();

	//printf("The return value: %d\n",GPU_filter(h_data,h_output,size,dimBlock,dimGrid,width,height));
	//time (&end);
	printf("%4.8f\n\n",(float)( std::clock() - start ));
  	 diff = ( std::clock() - start ) / (double)CLOCKS_PER_SEC;
	

	//if(i==1024)
	{
	//printElements(h_output,width,height);
	}

	printf ("Size :%d, Execution Time :%4.8f\n",i,diff );

	free(h_data);
	free(h_output);
	
	i=i*2-1;
	
	}

	return 0;

}
void initialize(int* elements,int width,int height)
{
	int n = width*height;
	int rowCount=0;
	for(int i=0;i<n;i++)
	{
		elements[i]= 0;
		if(rowCount>=8&&rowCount<=(height-8))
		{		
			if((rowCount*width+8)<=i&&(rowCount*width+(height-8))>=i)
			{
				elements[i]=255;				
			}
		}

		if((i+1)%width==0)
		{		
			rowCount++;
		}		
	}
}


void printElements(int* elements,int width,int height)
{
	int n = width*height;
	for(int i=0;i<n;i++)
	{	
		if(i%width==0)
		{
			printf("\n");
		}
		printf("%d ",elements[i]);
		
	}
	printf("\n");
}*/